
/* Parser-tokenizer link implementation */
#include "pgenheaders.h"
#include "tokenizer.h"
#include "node.h"
#include "grammar.h"
#include "parser.h"
#include "parsetok.h"
#include "errcode.h"
#include "graminit.h"


/* Tokenizer implementation */

#include "Python.h"

#include <cuda_runtime.h>
#include <cutil_inline.h>
#include "cuda_tokenizer.h"
#include "cuda_tokenizer_util.h"

/* cuda_strchr() cuda_strcpy() cuda_strstr()... for __device__ originally from <ctype.h> */
__device__
char *
cuda_strchr(char *string, int ch) {
    while (*string && *string != (char) ch)
        string++;

    if (*string == (char) ch)
        return ((char *) string);
    return (NULL);
}

__device__
char *
cuda_strcpy(char * strDest, char * strSrc) {
    char * strDestCopy = strDest;
    if ((strDest == NULL) || (strSrc == NULL))
        return NULL;
    while ((*strDest++ = *strSrc++) != '\0');

    return strDestCopy;
}

__device__
char *
cuda_strstr(char *s1, char *s2) {
    int n;

    if (*s2) {
        while (*s1) {
            for (n = 0; *(s1 + n) == *(s2 + n); n++) {
                if (!*(s2 + n + 1))
                    return (char *) s1;
            }
            s1++;
        }

        return NULL;
    } else
        return (char *) s1;
}

__device__
int cuda_atoi(char *s) {
    char *p = s;
    char c;
    int i = 0;

    while (c = *p++) {
        if (c >= '0' && c <= '9') {
            i = i * 10 + (c - '0');
        } else
            return -1; // Invalid string
    }
    return i;
}

__device__
int cuda_strlen(char *str) {
    //assert(str != NULL);
    int len = 0;

    while ((*str++) != '\0')
        len++;

    return len;
}

/* Back-up one character for GPU device */
__device__
void
cuda_tok_backup(struct tok_state *tok, int c) {
    if (c != EOF) {
        if (*tok->cur != c)
            *tok->cur = c;
    }
}

/* Get next char, updating state; error code goes into tok->done */
__device__ int
cuda_tok_nextc(char *d_buffer, struct tok_state *tok, const int *d_linetable, const int line_idx) { /*** @TODO #if 0 will be removed later */

    for (;;) {
        if (tok->cur != tok->inp) {
            return Py_CHARMASK(*tok->cur++); /* Fast path */
        }
        if (tok->done != E_OK)
            return EOF;
        if (tok->fp == NULL) {
            char *end = cuda_strchr(tok->inp, '\n');
            if (end != NULL)
                end++;
            else {
                end = cuda_strchr(tok->inp, '\0');
                if (end == tok->inp) {
                    tok->done = E_EOF;
                    return EOF;
                }
            }
            if (tok->start == NULL)
                tok->buf = tok->cur;
            tok->line_start = tok->cur;
            tok->lineno++;
            tok->inp = end;
            return Py_CHARMASK(*tok->cur++);
        }
        //        if (tok->prompt != NULL) {
        //            char *newtok = PyOS_Readline(stdin, stdout, tok->prompt);
        //            if (tok->nextprompt != NULL)
        //                tok->prompt = tok->nextprompt;
        //            if (newtok == NULL)
        //                tok->done = E_INTR;
        //            else if (*newtok == '\0') {
        //                PyMem_FREE(newtok);
        //                tok->done = E_EOF;
        //            }
        ////    #if !defined(PGEN) && defined(Py_USING_UNICODE)
        ////    else if (tok_stdin_decode(tok, &newtok) != 0)
        ////    PyMem_FREE(newtok);
        ////    #endif
        //            else if (tok->start != NULL) {
        //                size_t start = tok->start - tok->buf;
        //                size_t oldlen = tok->cur - tok->buf;
        //                size_t newlen = oldlen + cuda_strlen(newtok);
        //                char *buf = tok->buf;
        //                buf = (char *) PyMem_REALLOC(buf, newlen + 1);
        //                tok->lineno++;
        //                if (buf == NULL) {
        //                    PyMem_FREE(tok->buf);
        //                    tok->buf = NULL;
        //                    PyMem_FREE(newtok);
        //                    tok->done = E_NOMEM;
        //                    return EOF;
        //                }
        //                tok->buf = buf;
        //                tok->cur = tok->buf + oldlen;
        //                tok->line_start = tok->cur;
        //                cuda_strcpy(tok->buf + oldlen, newtok);
        //                PyMem_FREE(newtok);
        //                tok->inp = tok->buf + newlen;
        //                tok->end = tok->inp + 1;
        //                tok->start = tok->buf + start;
        //            } else {
        //                tok->lineno++;
        //                if (tok->buf != NULL)
        //                    PyMem_FREE(tok->buf);
        //                tok->buf = newtok;
        //                tok->line_start = tok->buf;
        //                tok->cur = tok->buf;
        //                tok->line_start = tok->buf;
        //                tok->inp = cuda_strchr(tok->buf, '\0');
        //                tok->end = tok->inp + 1;
        //            }
        //        } else 
        {
            int done = 0;
            Py_ssize_t cur = 0;
            char *pt;
            if (tok->start == NULL) {
                /*   if (tok->buf == NULL) {
                tok->buf = (char *)
                PyMem_MALLOC(BUFSIZ);
                if (tok->buf == NULL) {
                tok->done = E_NOMEM;
                return EOF;
                }
                tok->end = tok->buf + BUFSIZ;
                }
                if (cuda_decoding_fgets(tok->buf, (int)(tok->end - tok->buf),
                tok,d_linetable, line_idx) == NULL) {
                tok->done = E_EOF;
                done = 1;
                }
                else {*/
                tok->buf = d_buffer + d_linetable[line_idx];
                tok->done = E_OK;
                tok->inp = cuda_strchr(tok->buf, '\n');
                done = tok->inp[-1] == '\n';
                // }
            }/** wei
    else {
    cur = tok->cur - tok->buf;
    if (decoding_feof(tok)) {
    tok->done = E_EOF;
    done = 1;
    }
    else
    tok->done = E_OK;
    }*/
            tok->lineno++;
            /* Read until '\n' or EOF */
            while (!done) {
                Py_ssize_t curstart = tok->start == NULL ? -1 :
                        tok->start - tok->buf;
                Py_ssize_t curvalid = tok->inp - tok->buf;
                Py_ssize_t newsize = curvalid + BUFSIZ;
                char *newbuf = tok->buf;
                //newbuf = (char *)malloc(newsize);

                if (newbuf == NULL) {
                    tok->done = E_NOMEM;
                    tok->cur = tok->inp;
                    return EOF;
                }
                tok->buf = newbuf;
                tok->inp = tok->buf + curvalid;
                tok->end = tok->buf + newsize;
                tok->start = curstart < 0 ? NULL :
                        tok->buf + curstart;
                //#if 0                         
                //                if (cuda_decoding_fgets(tok->inp,
                //                               (int)(tok->end - tok->inp),
                //                               tok) == NULL) {
                //                    /* Break out early on decoding
                //                       errors, as tok->buf will be NULL
                //                     */
                //                    if (tok->decoding_erred)
                //                        return EOF;
                //                    /* Last line does not end in \n,
                //                       fake one */
                //                    cuda_strcpy(tok->inp, "\n");
                //                }
                //#endif
                tok->inp = d_buffer + d_linetable[line_idx];
                if (tok->decoding_erred)
                    return EOF;
                cuda_strcpy(tok->inp, "\n");

                tok->inp = cuda_strchr(tok->inp, '\0');
                done = tok->inp[-1] == '\n';
            }
            if (tok->buf != NULL) {
                tok->cur = tok->buf + cur;
                tok->line_start = tok->cur;
                /* replace "\r\n" with "\n" */
                /* For Mac leave the \r, giving a syntax error */
                pt = tok->inp - 2;
                if (pt >= tok->buf && *pt == '\r') {
                    *pt++ = '\n';
                    *pt = '\0';
                    tok->inp = pt;
                }
            }
            //        }
            if (tok->done != E_OK) {
                //            if (tok->prompt != NULL)
                //                PySys_WriteStderr("\n");
                tok->cur = tok->inp;
                return EOF;
            }
        }
    }
    /*NOTREACHED*/
}

/* Return the token corresponding to a single character */
__device__
int cuda_PyToken_OneChar(int c) {
    switch (c) {
        case '(': return LPAR;
        case ')': return RPAR;
        case '[': return LSQB;
        case ']': return RSQB;
        case ':': return COLON;
        case ',': return COMMA;
        case ';': return SEMI;
        case '+': return PLUS;
        case '-': return MINUS;
        case '*': return STAR;
        case '/': return SLASH;
        case '|': return VBAR;
        case '&': return AMPER;
        case '<': return LESS;
        case '>': return GREATER;
        case '=': return EQUAL;
        case '.': return DOT;
        case '%': return PERCENT;
        case '`': return BACKQUOTE;
        case '{': return LBRACE;
        case '}': return RBRACE;
        case '^': return CIRCUMFLEX;
        case '~': return TILDE;
        case '@': return AT;
        default: return OP;
    }
}

__device__
int
cuda_PyToken_TwoChars(int c1, int c2) {
    switch (c1) {
        case '=':
            switch (c2) {
                case '=': return EQEQUAL;
            }
            break;
        case '!':
            switch (c2) {
                case '=': return NOTEQUAL;
            }
            break;
        case '<':
            switch (c2) {
                case '>': return NOTEQUAL;
                case '=': return LESSEQUAL;
                case '<': return LEFTSHIFT;
            }
            break;
        case '>':
            switch (c2) {
                case '=': return GREATEREQUAL;
                case '>': return RIGHTSHIFT;
            }
            break;
        case '+':
            switch (c2) {
                case '=': return PLUSEQUAL;
            }
            break;
        case '-':
            switch (c2) {
                case '=': return MINEQUAL;
            }
            break;
        case '*':
            switch (c2) {
                case '*': return DOUBLESTAR;
                case '=': return STAREQUAL;
            }
            break;
        case '/':
            switch (c2) {
                case '/': return DOUBLESLASH;
                case '=': return SLASHEQUAL;
            }
            break;
        case '|':
            switch (c2) {
                case '=': return VBAREQUAL;
            }
            break;
        case '%':
            switch (c2) {
                case '=': return PERCENTEQUAL;
            }
            break;
        case '&':
            switch (c2) {
                case '=': return AMPEREQUAL;
            }
            break;
        case '^':
            switch (c2) {
                case '=': return CIRCUMFLEXEQUAL;
            }
            break;
    }
    return OP;
}

__device__
int
cuda_PyToken_ThreeChars(int c1, int c2, int c3) {
    switch (c1) {
        case '<':
            switch (c2) {
                case '<':
                    switch (c3) {
                        case '=':
                            return LEFTSHIFTEQUAL;
                    }
                    break;
            }
            break;
        case '>':
            switch (c2) {
                case '>':
                    switch (c3) {
                        case '=':
                            return RIGHTSHIFTEQUAL;
                    }
                    break;
            }
            break;
        case '*':
            switch (c2) {
                case '*':
                    switch (c3) {
                        case '=':
                            return DOUBLESTAREQUAL;
                    }
                    break;
            }
            break;
        case '/':
            switch (c2) {
                case '/':
                    switch (c3) {
                        case '=':
                            return DOUBLESLASHEQUAL;
                    }
                    break;
            }
            break;
    }
    return OP;
}

__device__ int
cuda_indenterror(struct tok_state *tok) {
    if (tok->alterror) {
        tok->done = E_TABSPACE;
        tok->cur = tok->inp;
        return 1;
    }
    if (tok->altwarning) {
        /*PySys_WriteStderr("%s: inconsistent use of tabs and spaces "
                          "in indentation\n", tok->filename);*/

        tok->altwarning = 0;
    }
    return 0;
}

/* Get next token, after space stripping etc. */
__device__ void
cuda_tok_get(char *d_buffer, struct tok_state *tok, int *p_start,
        int *p_end, int *result, const int* d_linetable, const int line_idx) {
    __shared__ int c;
    int blankline;
    
    *p_start = *p_end = 0;
nextline:
    tok->start = NULL;
    blankline = 0;

    /* Get indentation level */
    if (tok->atbol) {
        int col = 0;
        int altcol = 0;
        tok->atbol = 0;
        for (;;) {
            c = cuda_tok_nextc(d_buffer, tok, d_linetable, line_idx);
            if (c == ' ')
                col++, altcol++;
            else if (c == '\t') {
                col = (col / tok->tabsize + 1) * tok->tabsize;
                altcol = (altcol / tok->alttabsize + 1)
                        * tok->alttabsize;
            } else if (c == '\014') /* Control-L (formfeed) */
                col = altcol = 0; /* For Emacs users */
            else
                break;
        }
        cuda_tok_backup(tok, c);
        if (c == '#' || c == '\n') {
            /* Lines with only whitespace and/or comments
               shouldn't affect the indentation and are
               not passed to the parser as NEWLINE tokens,
               except *totally* empty lines in interactive
               mode, which signal the end of a command group. */
            if (col == 0 && c == '\n' && tok->prompt != NULL)
                blankline = 0; /* Let it through */
            else
                blankline = 1; /* Ignore completely */
            /* We can't jump back right here since we still
               may need to skip to the end of a comment */
        }
        if (!blankline && tok->level == 0) {
            if (col == tok->indstack[tok->indent]) {
                /* No change */
                if (altcol != tok->altindstack[tok->indent]) {
                    if (cuda_indenterror(tok)) {
                        (*result) = ERRORTOKEN;
                        return;
                    }
                }
            } else if (col > tok->indstack[tok->indent]) {
                /* Indent -- always one */
                if (tok->indent + 1 >= MAXINDENT) {
                    tok->done = E_TOODEEP;
                    tok->cur = tok->inp;
                    (*result) = ERRORTOKEN;
                    return;
                }
                if (altcol <= tok->altindstack[tok->indent]) {
                    if (cuda_indenterror(tok)) {
                        (*result) = ERRORTOKEN;
                        return;
                    }
                }
                tok->pendin++;
                tok->indstack[++tok->indent] = col;
                tok->altindstack[tok->indent] = altcol;
            } else /* col < tok->indstack[tok->indent] */ {
                /* Dedent -- any number, must be consistent */
                while (tok->indent > 0 &&
                        col < tok->indstack[tok->indent]) {
                    tok->pendin--;
                    tok->indent--;
                }
                if (col != tok->indstack[tok->indent]) {
                    tok->done = E_DEDENT;
                    tok->cur = tok->inp;
                    (*result) = ERRORTOKEN;
                    return;
                }
                if (altcol != tok->altindstack[tok->indent]) {
                    if (cuda_indenterror(tok)) {
                        (*result) = ERRORTOKEN;
                        return;
                    }
                }
            }
        }
    }

    tok->start = tok->cur;

    /* Return pending indents/dedents */
    if (tok->pendin != 0) {
        if (tok->pendin < 0) {
            tok->pendin++;
            (*result) = DEDENT;
            return;
        } else {
            tok->pendin--;
            (*result) = INDENT;
            return;
        }
    }

again:
    tok->start = NULL;
    /* Skip spaces */
    do {
        c = cuda_tok_nextc(d_buffer, tok, d_linetable, line_idx);
    } while (c == ' ' || c == '\t' || c == '\014');

    /* Set start of current token */
    tok->start = tok->cur - 1;

    /* Skip comment, while looking for tab-setting magic */
    if (c == '#') {
        /**static */char *tabforms[] = {
            "tab-width:", /* Emacs */
            ":tabstop=", /* vim, full form */
            ":ts=", /* vim, abbreviated form */
            "set tabsize=", /* will vi never die? */
            /* more templates can be added here to support other editors */
        };
        char cbuf[80];
        char *tp, **cp;
        tp = cbuf;
        do {
            *tp++ = c = cuda_tok_nextc(d_buffer, tok, d_linetable, line_idx);
        } while (c != EOF && c != '\n' &&
                (size_t) (tp - cbuf + 1) < sizeof (cbuf));
        *tp = '\0';
        for (cp = tabforms;
                cp < tabforms + sizeof (tabforms) / sizeof (tabforms[0]);
                cp++) {
            if ((tp = cuda_strstr(cbuf, *cp))) {
                int newsize = cuda_atoi(tp + cuda_strlen(*cp));

                if (newsize >= 1 && newsize <= 40) {
                    tok->tabsize = newsize;
                    /*if (Py_VerboseFlag)
                        PySys_WriteStderr(
                        "Tab size set to %d\n",
                        newsize);*/
                }
            }
        }
        while (c != EOF && c != '\n')
            c = cuda_tok_nextc(d_buffer, tok, d_linetable, line_idx);
    }

    /* Check for EOF and errors now */
    if (c == EOF) {
        (*result) = (tok->done == E_EOF) ? ENDMARKER : ERRORTOKEN;
        return;
    }

    /* Identifier (most frequent token!) */
    if (Cuda_ISALPHA(c) || c == '_') {
        /* Process r"", u"" and ur"" */
        switch (c) {
            case 'b':
            case 'B':
                c = cuda_tok_nextc(d_buffer, tok, d_linetable, line_idx);
                if (c == 'r' || c == 'R')
                    c = cuda_tok_nextc(d_buffer, tok, d_linetable, line_idx);
                if (c == '"' || c == '\'')
                    goto letter_quote;
                break;
            case 'r':
            case 'R':
                c = cuda_tok_nextc(d_buffer, tok, d_linetable, line_idx);
                if (c == '"' || c == '\'')
                    goto letter_quote;
                break;
            case 'u':
            case 'U':
                c = cuda_tok_nextc(d_buffer, tok, d_linetable, line_idx);
                if (c == 'r' || c == 'R')
                    c = cuda_tok_nextc(d_buffer, tok, d_linetable, line_idx);
                if (c == '"' || c == '\'')
                    goto letter_quote;
                break;
        }
        while (c != EOF && (Cuda_ISALNUM(c) || c == '_')) {
            c = cuda_tok_nextc(d_buffer, tok, d_linetable, line_idx);
        }
        cuda_tok_backup(tok, c);
        *p_start = cuda_atoi(tok->start);
        *p_end = cuda_atoi(tok->cur);
        (*result) = NAME;
        return;
    }

    /* Newline */
    if (c == '\n') {
        tok->atbol = 1;
        if (blankline || tok->level > 0)
            goto nextline;
        *p_start = cuda_atoi(tok->start);
        *p_end = cuda_atoi(tok->cur) - 1; /* Leave '\n' out of the string */
        tok->cont_line = 0;
        (*result) = NEWLINE;
        return;
    }

    /* Period or number starting with period? */
    if (c == '.') {
        c = cuda_tok_nextc(d_buffer, tok, d_linetable, line_idx);
        if (Cuda_ISDIGIT(c)) {
            goto fraction;
        } else {
            cuda_tok_backup(tok, c);
            *p_start = cuda_atoi(tok->start);
            *p_end = cuda_atoi(tok->cur);
            (*result) = DOT;
            return;
        }
    }

    /* Number */
    if (Cuda_ISDIGIT(c)) {
        if (c == '0') {
            /* Hex, octal or binary -- maybe. */
            c = cuda_tok_nextc(d_buffer, tok, d_linetable, line_idx);
            if (c == '.')
                goto fraction;
#ifndef WITHOUT_COMPLEX
            if (c == 'j' || c == 'J')
                goto imaginary;
#endif
            if (c == 'x' || c == 'X') {

                /* Hex */
                c = cuda_tok_nextc(d_buffer, tok, d_linetable, line_idx);
                if (!Cuda_ISXDIGIT(c)) {
                    tok->done = E_TOKEN;
                    cuda_tok_backup(tok, c);
                    (*result) = ERRORTOKEN;
                    return;
                }
                do {
                    c = cuda_tok_nextc(d_buffer, tok, d_linetable, line_idx);
                } while (Cuda_ISXDIGIT(c));
            } else if (c == 'o' || c == 'O') {
                /* Octal */
                c = cuda_tok_nextc(d_buffer, tok, d_linetable, line_idx);
                if (c < '0' || c >= '8') {
                    tok->done = E_TOKEN;
                    cuda_tok_backup(tok, c);
                    (*result) = ERRORTOKEN;
                    return;
                }
                do {
                    c = cuda_tok_nextc(d_buffer, tok, d_linetable, line_idx);
                } while ('0' <= c && c < '8');
            } else if (c == 'b' || c == 'B') {
                /* Binary */
                c = cuda_tok_nextc(d_buffer, tok, d_linetable, line_idx);
                if (c != '0' && c != '1') {
                    tok->done = E_TOKEN;
                    cuda_tok_backup(tok, c);
                    (*result) = ERRORTOKEN;
                    return;
                }
                do {
                    c = cuda_tok_nextc(d_buffer, tok, d_linetable, line_idx);
                } while (c == '0' || c == '1');
            } else {
                int found_decimal = 0;
                /* Octal; c is first char of it */
                /* There's no 'isoctdigit' macro, sigh */
                while ('0' <= c && c < '8') {
                    c = cuda_tok_nextc(d_buffer, tok, d_linetable, line_idx);
                }
                if (Cuda_ISDIGIT(c)) {
                    found_decimal = 1;
                    do {
                        c = cuda_tok_nextc(d_buffer, tok, d_linetable, line_idx);
                    } while (Cuda_ISDIGIT(c));
                }
                if (c == '.')
                    goto fraction;
                else if (c == 'e' || c == 'E')
                    goto exponent;
#ifndef WITHOUT_COMPLEX
                else if (c == 'j' || c == 'J')
                    goto imaginary;
#endif
                else if (found_decimal) {
                    tok->done = E_TOKEN;
                    cuda_tok_backup(tok, c);
                    (*result) = ERRORTOKEN;
                    return;
                }
            }
            if (c == 'l' || c == 'L')
                c = cuda_tok_nextc(d_buffer, tok, d_linetable, line_idx);
        } else {
            /* Decimal */
            do {
                c = cuda_tok_nextc(d_buffer, tok, d_linetable, line_idx);
            } while (Cuda_ISDIGIT(c));
            if (c == 'l' || c == 'L')
                c = cuda_tok_nextc(d_buffer, tok, d_linetable, line_idx);
            else {
                /* Accept floating point numbers. */
                if (c == '.') {
fraction:
                    /* Fraction */
                    do {
                        c = cuda_tok_nextc(d_buffer, tok, d_linetable, line_idx);
                    } while (Cuda_ISDIGIT(c));
                }
                if (c == 'e' || c == 'E') {
exponent:
                    /* Exponent part */
                    c = cuda_tok_nextc(d_buffer, tok, d_linetable, line_idx);
                    if (c == '+' || c == '-')
                        c = cuda_tok_nextc(d_buffer, tok, d_linetable, line_idx);
                    if (!Cuda_ISDIGIT(c)) {
                        tok->done = E_TOKEN;
                        cuda_tok_backup(tok, c);
                        (*result) = ERRORTOKEN;
                        return;
                    }
                    do {
                        c = cuda_tok_nextc(d_buffer, tok, d_linetable, line_idx);
                    } while (Cuda_ISDIGIT(c));
                }
#ifndef WITHOUT_COMPLEX
                if (c == 'j' || c == 'J')
                    /* Imaginary part */
                    imaginary :
                        c = cuda_tok_nextc(d_buffer, tok, d_linetable, line_idx);
#endif
            }
        }
        cuda_tok_backup(tok, c);
        *p_start = cuda_atoi(tok->start);
        *p_end = cuda_atoi(tok->cur);
        (*result) = NUMBER;
        return;
    }

letter_quote:
    /* String */
    if (c == '\'' || c == '"') {
        Py_ssize_t quote2 = tok->cur - tok->start + 1;
        int quote = c;
        int triple = 0;
        int tripcount = 0;
        for (;;) {
            c = cuda_tok_nextc(d_buffer, tok, d_linetable, line_idx);
            if (c == '\n') {
                if (!triple) {
                    tok->done = E_EOLS;
                    cuda_tok_backup(tok, c);
                    (*result) = ERRORTOKEN;
                    return;
                }
                tripcount = 0;
                tok->cont_line = 1; /* multiline string. */
            } else if (c == EOF) {
                if (triple)
                    tok->done = E_EOFS;
                else
                    tok->done = E_EOLS;
                tok->cur = tok->inp;
                (*result) = ERRORTOKEN;
                return;
            } else if (c == quote) {
                tripcount++;
                if (tok->cur - tok->start == quote2) {
                    c = cuda_tok_nextc(d_buffer, tok, d_linetable, line_idx);
                    if (c == quote) {
                        triple = 1;
                        tripcount = 0;
                        continue;
                    }
                    cuda_tok_backup(tok, c);
                }
                if (!triple || tripcount == 3)
                    break;
            } else if (c == '\\') {
                tripcount = 0;
                c = cuda_tok_nextc(d_buffer, tok, d_linetable, line_idx);
                if (c == EOF) {
                    tok->done = E_EOLS;
                    tok->cur = tok->inp;
                    (*result) = ERRORTOKEN;
                    return;
                }
            } else
                tripcount = 0;
        }
        *p_start = cuda_atoi(tok->start);
        *p_end = cuda_atoi(tok->cur);
        (*result) = STRING;
        return;
    }

    /* Line continuation */
    if (c == '\\') {
        c = cuda_tok_nextc(d_buffer, tok, d_linetable, line_idx);
        if (c != '\n') {
            tok->done = E_LINECONT;
            tok->cur = tok->inp;
            (*result) = ERRORTOKEN;
            return;
        }
        tok->cont_line = 1;
        goto again; /* Read next line */
    }

    /* Check for two-character token */
    {
        int c2 = cuda_tok_nextc(d_buffer, tok, d_linetable, line_idx);
        int token = cuda_PyToken_TwoChars(c, c2);
#ifndef PGEN
        /**if (Py_Py3kWarningFlag && token == NOTEQUAL && c == '<') {
            if (PyErr_WarnExplicit(PyExc_DeprecationWarning,
                                   "<> not supported in 3.x; use !=",
                                   tok->filename, tok->lineno,
                                   NULL, NULL)) {
                (*result) = ERRORTOKEN;
                return;
            }
        }*/
#endif
        if (token != OP) {
            int c3 = cuda_tok_nextc(d_buffer, tok, d_linetable, line_idx);
            int token3 = cuda_PyToken_ThreeChars(c, c2, c3);
            if (token3 != OP) {
                token = token3;
            } else {
                cuda_tok_backup(tok, c3);
            }
            *p_start = cuda_atoi(tok->start);
            *p_end = cuda_atoi(tok->cur);
            (*result) = token;
            return;
        }
        cuda_tok_backup(tok, c2);
    }

    /* Keep track of parentheses nesting level */
    switch (c) {
        case '(':
        case '[':
        case '{':
            tok->level++;
            break;
        case ')':
        case ']':
        case '}':
            tok->level--;
            break;
    }

    /* Punctuation character */
    *p_start = cuda_atoi(tok->start);
    *p_end = cuda_atoi(tok->cur);
    (*result) = cuda_PyToken_OneChar(c);
    return;
}

__global__ void
cuda_PyTokenizer_Get(char *d_buffer, struct cuda_tok_state *d_cuda_tok,
        struct tok_state *d_tok, const int *d_linetable, const int newline_num) {

    const int bid = blockIdx.x;
    const int tid = threadIdx.x;
    const int threadid = bid * BLOCK_NUM + tid;
    int quotient = newline_num / (BLOCK_NUM * THREAD_NUM);
    const int remainder = newline_num % (BLOCK_NUM * THREAD_NUM);
    int bufIndex;
    int idx = bid * 4096 + tid * 256;

    if (threadid >= 0 && threadid < remainder) {
        bufIndex = (quotient + 1) * threadid;
        quotient++;
    }
    else {
        bufIndex = (quotient + 1) * remainder + quotient * (threadid - remainder);
    }

    int i;
    int count = 0;

    for (i = 0; i < quotient; i++) {

        int *p_start, *p_end;
        int result = 0;

        do {
            count++;

            //////////////////////
            char * d_line_buffer = d_buffer + d_linetable[bufIndex + i];
            //////////////////////

//            cuda_tok_get(d_line_buffer, d_tok, p_start, p_end, &result,
//                    d_linetable, bufIndex + i);

            //////////////////////
//            if ((*p_end - *p_start) >= (d_linetable[bufIndex + i + 1] - d_linetable[bufIndex + i]))
//                continue;
            ///////////////////////

            if (d_tok->decoding_erred) {
                result = ERRORTOKEN;
                d_tok->done = E_DECODE;
            }
//            d_cuda_tok[idx + count].type = result;
//            d_cuda_tok[idx + count].offset_a = *p_start - cuda_atoi(d_buffer);
//            d_cuda_tok[idx + count].offset_b = *p_end - cuda_atoi(d_buffer);
//            d_cuda_tok[idx + count].offset_line_start = d_tok->line_start - d_buffer;
//            d_cuda_tok[idx + count].lineno = d_tok->lineno;
//            d_cuda_tok[idx + count].done = d_tok->done;
//            d_cuda_tok[idx + count].indent = d_tok->indent;
            
            result = NEWLINE;
        } while (result != NEWLINE);
    }

    __syncthreads();

    if (bid == 0 && tid == 0)
        d_cuda_tok[0].type = 111;

    //return result;
    /* */
}


/* Parse input coming from the given tokenizer structure.
   Return error code. */

node *
cuda_parsetok(struct tok_state *tok, grammar *g, int start, perrdetail *err_ret,
        int *flags) {
    struct timeval start_tokenizer, stop_tokenizer;
    struct timeval start_host2device, stop_host2device;
    struct timeval start_device2host, stop_device2host;
    struct timeval start_file, stop_file;
    
    gettimeofday(&start_tokenizer, NULL);
    parser_state *ps;
    node *n;
    int started = 0, handling_import = 0, handling_with = 0;

    if ((ps = PyParser_New(g, start)) == NULL) {
        fprintf(stderr, "no mem for new parser\n");
        err_ret->error = E_NOMEM;
        PyTokenizer_Free(tok);
        return NULL;
    }

#ifdef PY_PARSER_REQUIRES_FUTURE_KEYWORD
    if (*flags & PyPARSE_PRINT_IS_FUNCTION) {
        ps->p_flags |= CO_FUTURE_PRINT_FUNCTION;
    }
    if (*flags & PyPARSE_UNICODE_LITERALS) {
        ps->p_flags |= CO_FUTURE_UNICODE_LITERALS;
    }

#endif

    char *d_buffer, h_buffer[FILE_SIZE];
    struct cuda_tok_state *d_cuda_tok, *h_cuda_tok = NULL;
    struct tok_state *d_tok, *h_tok = NULL;


    // 256KB
    cutilSafeCall(cudaMalloc((void **) &d_buffer, sizeof (char) * FILE_SIZE));
    // 14MB
    cutilSafeCall(cudaMalloc((void **) &d_cuda_tok, DATA_SIZE * sizeof (struct cuda_tok_state)));
    // 8KB = 2K * 4
    cutilSafeCall(cudaMalloc((void **) &d_tok, sizeof (struct tok_state)));


    h_cuda_tok = (struct cuda_tok_state *) PyMem_MALLOC(DATA_SIZE * sizeof (struct cuda_tok_state));
    h_tok = (struct tok_state *) PyMem_MALLOC(sizeof (struct tok_state));

    
    size_t linecap = FILE_SIZE;
    ssize_t linelen;
    int atline;
    int *h_linetable = (int *) PyMem_MALLOC(sizeof (int) * 8192);
    int *d_linetable;

    h_linetable[0] = 0;

    char *h_buffer_tmp_p = h_buffer;
    
    gettimeofday(&start_file, NULL);
    /* Read file from tok->fp to h_buffer. */
    if ((tok->fp = fopen(tok->filename, "r")) == NULL) {
        fprintf(stderr, "Open failed! %s\n", tok->filename);
        exit(1);
    }

    for (atline = 1; (linelen = getline(&h_buffer_tmp_p, &linecap, tok->fp)) > 0; atline++) {
        h_buffer_tmp_p += linelen;
        h_linetable[atline] = h_linetable[atline - 1] + linelen;
    }

    fclose(tok->fp);
    gettimeofday(&stop_file, NULL);
    //h_buffer_tmp_p[h_linetable[atline - 1]] = 0xFF;

    cutilSafeCall(cudaMalloc((void **) &d_linetable, sizeof (int) * 8192));
    
    gettimeofday(&start_host2device, NULL);
    cutilSafeCall(cudaMemcpy(d_linetable, h_linetable, sizeof (int) * 8192,
            cudaMemcpyHostToDevice));

    cutilSafeCall(cudaMemcpy(d_tok, tok, sizeof (struct tok_state),
            cudaMemcpyHostToDevice));

    cutilSafeCall(cudaMemcpy(d_buffer, h_buffer, sizeof (char) * (h_linetable[atline - 1] + 1),
            cudaMemcpyHostToDevice));
    gettimeofday(&stop_host2device, NULL);
    
    cuda_PyTokenizer_Get << <BLOCK_NUM, THREAD_NUM >> >(d_buffer, d_cuda_tok,
            d_tok, d_linetable, (atline - 1));

//    printf("%s\n", cudaGetErrorString(cudaGetLastError()));
//    printf("Sync: %s\n", cudaGetErrorString(cudaThreadSynchronize()));

    cutilDeviceSynchronize();

    // check if kernel execution generated and error
    cutilCheckMsg("CUDA Kernel execution failed");

    gettimeofday(&start_device2host, NULL);
    cutilSafeCall(cudaMemcpy(h_tok, d_tok, sizeof (struct tok_state),
            cudaMemcpyDeviceToHost));


    /* <-- printf("D: get into cuda_parsetok\n");
     * Cannot progress in the following statement. 
     * $h_token_num is invalid (== 0);
     */
    cutilSafeCall(cudaMemcpy(h_cuda_tok, d_cuda_tok,
            DATA_SIZE * sizeof (struct cuda_tok_state), cudaMemcpyDeviceToHost));
    gettimeofday(&stop_device2host, NULL);
    /* --> */

//    printf("h_cuda_tok[0].type = %d\nline_num = %d\n", h_cuda_tok[0].type, atline - 1);

    gettimeofday(&stop_tokenizer, NULL);

    int j;
    for (j = 0; j < (1000); j++) {
        char *a, *b;
        int type;
        size_t len;
        char *str;
        int col_offset;

        a = h_buffer + h_cuda_tok[j].offset_a;
        b = h_buffer + h_cuda_tok[j].offset_b;
        tok->line_start = h_buffer + h_cuda_tok[j].offset_line_start;
        tok->lineno = h_cuda_tok[j].lineno;
        tok->done = h_cuda_tok[j].done;

        type = h_cuda_tok[j].type;
        if (type == ERRORTOKEN) {
            err_ret->error = tok->done;
            break;
        }
        if (type == ENDMARKER && started) {
            type = NEWLINE; /* Add an extra newline */
            handling_with = handling_import = 0;
            started = 0;
            /* Add the right number of dedent tokens,
               except if a certain flag is given --
               codeop.py uses this. */
            if (tok->indent &&
                    !(*flags & PyPARSE_DONT_IMPLY_DEDENT)) {
                tok->pendin = -tok->indent;
                tok->indent = 0;
            }
        } else
            started = 1;
        len = b - a; /* XXX this may compute NULL - NULL */
//        printf("len = %ld\n", len);
        str = (char *) PyObject_MALLOC(len + 1);
        if (str == NULL) {
            fprintf(stderr, "no mem for next token\n");
            err_ret->error = E_NOMEM;
            break;
        }
        if (len > 0)
            strncpy(str, a, len);
        str[len] = '\0';

#ifdef PY_PARSER_REQUIRES_FUTURE_KEYWORD
#endif
        if (a >= tok->line_start)
            col_offset = a - tok->line_start;
        else
            col_offset = -1;

        if ((err_ret->error =
                PyParser_AddToken(ps, (int) type, str, tok->lineno, col_offset,
                &(err_ret->expected))) != E_OK) {
            if (err_ret->error != E_DONE) {
                PyObject_FREE(str);
                err_ret->token = type;
            }
            break;
        }
    }

    memcpy(tok, h_tok, sizeof (struct tok_state));

    if (err_ret->error == E_DONE) {
        n = ps->p_tree;
        ps->p_tree = NULL;
    } else
        n = NULL;

#ifdef PY_PARSER_REQUIRES_FUTURE_KEYWORD
    *flags = ps->p_flags;
#endif
    PyParser_Delete(ps);

    if (n == NULL) {
        if (tok->lineno <= 1 && tok->done == E_EOF)
            err_ret->error = E_EOF;
        err_ret->lineno = tok->lineno;
        if (tok->buf != NULL) {
            char *text = NULL;
            size_t len;
            assert(tok->cur - tok->buf < INT_MAX);
            err_ret->offset = (int) (tok->cur - tok->buf);
            len = tok->inp - tok->buf;
#ifdef Py_USING_UNICODE
            text = PyTokenizer_RestoreEncoding(tok, len, &err_ret->offset);

#endif
            if (text == NULL) {
                text = (char *) PyObject_MALLOC(len + 1);
                if (text != NULL) {
                    if (len > 0)
                        strncpy(text, tok->buf, len);
                    text[len] = '\0';
                }
            }
            err_ret->text = text;
        }
    } else if (tok->encoding != NULL) {
        /* 'nodes->n_str' uses PyObject_*, while 'tok->encoding' was
         * allocated using PyMem_
         */
        node* r = PyNode_New(encoding_decl);
        if (r)
            r->n_str = (char *) PyObject_MALLOC(strlen(tok->encoding) + 1);
        if (!r || !r->n_str) {
            err_ret->error = E_NOMEM;
            if (r)
                PyObject_FREE(r);
            n = NULL;
            goto done;
        }
        strcpy(r->n_str, tok->encoding);
        PyMem_FREE(tok->encoding);
        tok->encoding = NULL;
        r->n_nchildren = 1;
        r->n_child = n;
        n = r;
    }
    
done:
    PyTokenizer_Free(tok);

    cutilSafeCall(cudaFree(d_buffer));
    cutilSafeCall(cudaFree(d_cuda_tok));
    cutilSafeCall(cudaFree(d_tok));

    //    if (h_token_num != NULL)
    //        PyMem_FREE(h_token_num);


    if (h_linetable != NULL)
        PyMem_FREE(h_linetable);

    cutilSafeCall(cudaFree(d_linetable));

    cutilDeviceReset();

    
    // <Debug>
//    printf("D: get into cuda_parsetok\tn = %d\n", n);
    // </Debug>

    printf("stop_tokenizer - start_tokenizer = %ld (us)\n", 
            1000000 * (stop_tokenizer.tv_sec - start_tokenizer.tv_sec) 
            + stop_tokenizer.tv_usec - start_tokenizer.tv_usec);
    printf("stop_host2device - start_host2device = %ld (us)\n", 
            1000000 * (stop_host2device.tv_sec - start_host2device.tv_sec)
            + stop_host2device.tv_usec - start_host2device.tv_usec);
    printf("stop_device2host - start_device2host = %ld (us)\n", 
            1000000 * (stop_device2host.tv_sec - start_device2host.tv_sec)
            + stop_device2host.tv_usec - start_device2host.tv_usec);
    printf("stop_file - start_file = %ld (us)\n",
            1000000 * (stop_file.tv_sec - start_file.tv_sec)
            + stop_file.tv_usec - start_file.tv_usec);
    
    return n;
}
